//#pragma OPENCL EXTENSION cl_amd_printf : enable
//#define N_UNROLL  4  //unroll, future by mako

__kernel void zero_f_e_vir( 
    const int n_atoms,
    __global float4 *f_e,  //force, en
    __global float8 *vir   //virial
    ) 
{
    int iglb = get_global_id(0);
    int sglb = get_global_size(0);
    const float4 z4 = (float4)(0,0,0,0);  
    const float8 z8 = (float8)(0,0,0,0,0,0,0,0);  

    int i = iglb;
    while(i<n_atoms){
        f_e[i] = z4;
        vir[i] = z8;
        i += sglb;
    }
}

__kernel void zero_f_e( 
    const int n_atoms,
    __global float4 *f_e  //force, en
    ) 
{
    int iglb = get_global_id(0);
    int sglb = get_global_size(0);
    const float4 z4 = (float4)(0,0,0,0);  

    int i = iglb;
    while(i<n_atoms){
        f_e[i] = z4;
        i += sglb;
    }
}

__kernel void add_float4( 
    const int n_atoms,
    __global float4 *v1,   
    __global float4 *v2,  
    __global float4 *v_out
    ) 
{
    int iglb = get_global_id(0);
    int sglb = get_global_size(0);
    int i = iglb;
    while(i<n_atoms){
        v_out[i] = v1[i] + v2[i];
        i += sglb;
    }
}

__kernel void add_float8( 
    const int n_atoms,
    __global float8 *v1,   
    __global float8 *v2,  
    __global float8 *v_out
    ) 
{
    int iglb = get_global_id(0);
    int sglb = get_global_size(0);
    int i = iglb;
    while(i<n_atoms){
        v_out[i] = v1[i] + v2[i];
        i += sglb;
    }
}
